[KMCompiler][MetaX][Iluvatar][Hygon] Add linalg_lstsq MetaX, Iluvatar and Hygon backend support - #5390

Adds a MetaX override for linalg_lstsq. Every device-specific value is
applied by rebinding attributes on flag_gems.ops.linalg_lstsq at import,
so the generic operator is not modified at all -- the driver reads each
of them at call time, and duplicating its ~250-line compact-WY driver
would only drift the moment the generic side changes.

Three device facts, each measured on a C550 rather than inferred:

1. 64KB of shared memory per block. The tuned WY update config asks for
   73728 bytes in both dtypes, 8KB over, and 19 of the suite's cases died
   with OutOfResources. BLOCK_C is now derived from the limit Triton
   enforces (its driver's max_shared_mem, not torch's static figure, which
   would make an H20 look like a 48KB part). The target is HALF the limit
   so two blocks stay resident per SM: bc=16 ran the update in 0.0797 ms
   against bc=32 at 0.0977, because only one block fits at 40960 B.

2. float64 tl.dot is miscompiled. Relative error ~1.0 at every operand
   shape including 16x16x16, with and without input_precision="ieee",
   while the same contraction written as rank-1 updates is bit-exact.
   _wy_update_metax branches on the COMPUTE constexpr, so fp32 keeps the
   tl.dot path unchanged -- it is correct here (1.1e-07) and faster.

3. 4KB/thread of private (register-spill) memory, a driver setting, which
   the monolithic path's reduce kernel exceeds at BLOCK_NC=128 -- surfacing
   as the misleading "memory size or pointer value too large to fit in 32
   bit". num_warps=8 fits and is correct but costs 25x (14.9 vs 0.60 ms),
   and num_warps=16 exceeds threads-per-block, so _TALL_MAX_NC_F32 is
   capped at 64 and those shapes route to compact-WY, which has no ceiling
   by design.

The one change outside the vendor directory is in the test helper: the
reference is now built with to_reference(..., upcast=True). On backends
where to_reference keeps the tensor on the device, torch.linalg.lstsq is
the VENDOR's kernel rather than truth, and against a float64 CPU solve
this op was off by 5.2e-08 where MACA's own fp32 lstsq was off by 1.3e-04
-- so two cases were failing on the reference's error, not ours. Upcasting
is gated on fp64 support, so Ascend 910B keeps the fp32 reference.

tests/test_linalg_lstsq.py: 76 passed on a C550.

`to_reference(..., upcast=True)` fixed the reference's precision but left the
solve running on the vendor's device, which is the thing that was wrong with
it. On maca3810 there is no fp64 `torch.linalg.lstsq`, so the reference itself
raised `CUDA error: invalid device function` before the operator was ever
called -- the failing case, (16,4) fp32 with NC=5, takes the monolithic path
that this backend override does not alter at all.

Compute the reference on the CPU in float64 and copy the result back to the
device. That is the same reference on every backend, independent of vendor
kernel accuracy and availability. The dtype of the copy is gated on fp64
support, so devices without it (Ascend 910B) get an fp32 reference exactly as
before.

Verified on a MetaX C550: 76 passed.

test-op.sh runs a second pass with `--ref=cpu`, which sets TO_CPU. In that
mode gems_assert_close moves `res` to the CPU and asserts the reference is
ALREADY there, so unconditionally copying the new CPU reference back to the
device broke it -- the values matched exactly, only the device did not.

Place the reference where the comparison expects it: CPU under TO_CPU, on the
device otherwise.

Verified on a MetaX C550, both passes green: the default run and
`--ref=cpu --quick`.

float64 `tl.dot` does not compile on a BI-V150:

    test_linalg_lstsq_tall_blocked_fp64
    triton.compiler.errors.CompilationError: at 39:16:
        Wacc += tl.dot(tl.trans(Vb), Tb, input_precision="ieee")

Every other fp64 case passes, because the monolithic and blocked-TSQR paths
contain no `tl.dot` at all -- only compact-WY contracts with one, and all three
of the operator's `tl.dot` calls live in that single kernel. So the port is one
kernel: `_wy_update` with a float64 form expressing each contraction as P
rank-1 updates, selected by the COMPUTE constexpr so float32 keeps `tl.dot`.

Nothing else is rebound. Unlike MetaX there is no shared-memory pressure here,
and configuration not measured on the device does not belong in an override.

Three tests also built their reference by calling `torch.linalg.lstsq` on
DEVICE tensors, which runs the vendor's kernel rather than a trusted one.
Iluvatar's cuSOLVER shim has no float64 QR, so `square_wy_fp64` failed inside
its own reference with `cusolverDnDormqr_bufferSize` -- nothing an operator
override can reach. They now solve on the CPU like `_ref_and_gems`, so the
reference is the same on every backend. Nothing is skipped and no assertion is
dropped.

`near_singular` deliberately keeps torch on the device: it asserts gems returns
the same four-tuple contract torch does -- solution, residuals, rank and
singular_values -- so the device's own torch is the correct reference there.

Measured on a BI-V150: 74 passed / 2 failed before, 75 passed / 1 failed with
the override alone.

Both are harness portability, not operator behaviour, and both surfaced only
once the suite ran on more than one vendor.

square_wy_fp64 compares res against the reference with plain arithmetic rather
than through gems_assert_close, so under `--ref=cpu` -- where _cpu_ref
correctly leaves the reference on the CPU -- it mixed a cuda tensor with a cpu
one. Route it through utils.to_cpu, which is what that helper is for.

complex_fallback called torch.allclose on complex64 DEVICE tensors, which needs
an elementwise complex abs. Iluvatar's runtime compiler cannot build one:
`[IXRTC] nvrtcCompileProgram failed ... abs_kernel<std::complex<float>>`. The
comparison gains nothing from running on the device, so keep the reference on
the CPU and move the result to it.

`_require_dtype` asked only the operator, via a 4x2 probe. That probe routes to
the monolithic path, so it answers for one of four code paths and cannot speak
for the others: on an Iluvatar BI-V150 it succeeds while a 256x256 float64
solve returns NaN.

The device is unambiguous that float64 is not available:

  * flag_gems.runtime.device.support_fp64 is False
  * float64 tl.dot does not compile
  * torch.matmul: "gemm of double is not supported on CoreX"
  * cuSOLVER has neither Dormqr nor Dorgqr
  * every float64 copy warns "Limited support for torch.double"

So consult the declared flag first and skip float64 there, matching the
suite-wide convention (ALL_FLOAT_DTYPES already gates on fp64_is_supported).
The operator probe stays as a second gate, for backends that do have float64
but lack this particular kernel -- Ascend, which answers NotImplementedError.

The _wy_update override keeps its float64 branch: the compile error it removes
is real, it costs float32 nothing (COMPUTE is a constexpr), and it is what
makes compact-WY work if this device gains usable float64. Its module now says
plainly that CI does not exercise that branch.

The shared-memory fix cleared all 14 OutOfResources cases -- CI reached 49
passed -- and then hit a different wall:

    triton/language/semantic.py:1445
    AssertionError: Unsupported lhs dtype fp64
    CompilationError: at 39:16  (_wy_update's tl.dot)

That is Triton's own frontend allow-list, not the silicon: a local Hygon box
with a different Triton build runs the same float64 tests through `tl.dot` and
passes all 76. The CI runner also declares float64 supported, so the dtype gate
does not skip it, and _gems_supports probes with a 4x2 solve that routes to the
monolithic path and never reaches a `tl.dot` -- so nothing upstream of the
kernel can see this.

Carry the same float64 rank-1 form the Iluvatar override uses, branching on the
COMPUTE constexpr so float32 keeps `tl.dot`. The kernel is byte-identical to
that one, which is duplication worth naming: three backends now carry it, and a
single `USE_DOT` constexpr on the generic kernel would replace all three copies
with a one-line rebind each. That is the right shape once a fourth appears.

Note this makes float64 use rank-1 on Hygon boxes whose Triton would have
accepted `tl.dot`, trading some float64 throughput for one file that works on
both Triton builds.